/*
 * clutilits.cpp
 *
 *  Created on: 2016年2月25日
 *      Author: guyadong
 */
#include <utility>
#include <memory>
#include "cl_utilits.h"
#include "intrinsic_wrapper.h"
using namespace std;
namespace gdface {
#define case_value(v) case v:return #v;
#define check_add_flag(f) if(flags&f){if(str.size())str.append("|");str.append(#f);}


std::string cl_utilits::deviceTypeToString(cl_device_type type) {
	switch (type) {
	case_value(CL_DEVICE_TYPE_DEFAULT)
	case_value(CL_DEVICE_TYPE_CPU)
	case_value(CL_DEVICE_TYPE_GPU)
	case_value(CL_DEVICE_TYPE_ACCELERATOR)
	case_value(CL_DEVICE_TYPE_CUSTOM)
	case_value(CL_DEVICE_TYPE_ALL)
	default:
		return "Unknown device type";
	}
}
std::string cl_utilits::imageChannelOrderToString(cl_channel_order channelOrder) {
	switch (channelOrder) {
	case_value(CL_R)
	case_value(CL_A)
	case_value(CL_RG)
	case_value(CL_RA)
	case_value(CL_RGB)
	case_value(CL_RGBA)
	case_value(CL_BGRA)
	case_value(CL_ARGB)
	case_value(CL_INTENSITY)
	case_value(CL_LUMINANCE)
	case_value(CL_Rx)
	case_value(CL_RGx)
	case_value(CL_RGBx)
	default:
		return "Unknown image channel order";
	}
}

std::string cl_utilits::imageChannelDataTypeToString(cl_channel_type channelDataType) {
	switch (channelDataType) {
	case_value(CL_SNORM_INT8)
	case_value(CL_SNORM_INT16)
	case_value(CL_UNORM_INT8)
	case_value(CL_UNORM_INT16)
	case_value(CL_UNORM_SHORT_565)
	case_value(CL_UNORM_SHORT_555)
	case_value(CL_UNORM_INT_101010)
	case_value(CL_SIGNED_INT8)
	case_value(CL_SIGNED_INT16)
	case_value(CL_SIGNED_INT32)
	case_value(CL_UNSIGNED_INT8)
	case_value(CL_UNSIGNED_INT16)
	case_value(CL_UNSIGNED_INT32)
	case_value(CL_HALF_FLOAT)
	case_value(CL_FLOAT)
	case_value(CL_UNORM_INT24)
#ifdef CL_VERSION_2_1
	case_value(CL_UNORM_INT_101010_2)
#endif
	default:
		return "Unknown image channel  type";
	}
}
std::string cl_utilits::imageFormatToString(const cl::ImageFormat& format) {
	return imageChannelOrderToString(format.image_channel_order).append(",").append(
			imageChannelDataTypeToString(format.image_channel_data_type));
}

std::string cl_utilits::memFlagsToString(cl_mem_flags flags) {
	std::string str;
	check_add_flag(CL_MEM_READ_WRITE);
	check_add_flag(CL_MEM_WRITE_ONLY);
	check_add_flag(CL_MEM_READ_ONLY);
	check_add_flag(CL_MEM_USE_HOST_PTR);
	check_add_flag(CL_MEM_ALLOC_HOST_PTR);
	check_add_flag(CL_MEM_COPY_HOST_PTR);
	check_add_flag(CL_MEM_HOST_WRITE_ONLY);
	check_add_flag(CL_MEM_HOST_READ_ONLY);
	check_add_flag(CL_MEM_HOST_NO_ACCESS);
#if defined(CL_VERSION_2_0) || defined(CL_VERSION_2_1)
	check_add_flag(CL_MEM_SVM_FINE_GRAIN_BUFFER);
	check_add_flag(CL_MEM_SVM_ATOMICS);
	check_add_flag(CL_MEM_KERNEL_READ_AND_WRITE);
#endif
	return std::move(str);
}

std::string cl_utilits::memObjectTypeToString(cl_mem_object_type type) {
	switch (type) {
	case_value(CL_MEM_OBJECT_BUFFER)
	case_value(CL_MEM_OBJECT_IMAGE2D)
	case_value(CL_MEM_OBJECT_IMAGE3D)
	case_value(CL_MEM_OBJECT_IMAGE2D_ARRAY)
	case_value(CL_MEM_OBJECT_IMAGE1D)
	case_value(CL_MEM_OBJECT_IMAGE1D_ARRAY)
	case_value(CL_MEM_OBJECT_IMAGE1D_BUFFER)
#if defined(CL_VERSION_2_0) || defined(CL_VERSION_2_1)
	case_value(CL_MEM_OBJECT_PIPE)
#endif
	default:
		return "Unknown mem object type";
	}
}

std::string cl_utilits::errorNumberToString(cl_int errorNumber) {
	switch (errorNumber) {
	case_value(CL_SUCCESS)
	case_value(CL_DEVICE_NOT_FOUND)
	case_value(CL_DEVICE_NOT_AVAILABLE)
	case_value(CL_COMPILER_NOT_AVAILABLE)
	case_value(CL_MEM_OBJECT_ALLOCATION_FAILURE)
	case_value(CL_OUT_OF_RESOURCES)
	case_value(CL_OUT_OF_HOST_MEMORY)
	case_value(CL_PROFILING_INFO_NOT_AVAILABLE)
	case_value(CL_MEM_COPY_OVERLAP)
	case_value(CL_IMAGE_FORMAT_MISMATCH)
	case_value(CL_IMAGE_FORMAT_NOT_SUPPORTED)
	case_value(CL_BUILD_PROGRAM_FAILURE)
	case_value(CL_MAP_FAILURE)
	case_value(CL_MISALIGNED_SUB_BUFFER_OFFSET)
	case_value(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST)
	case_value(CL_COMPILE_PROGRAM_FAILURE)
	case_value(CL_LINKER_NOT_AVAILABLE)
	case_value(CL_LINK_PROGRAM_FAILURE)
	case_value(CL_DEVICE_PARTITION_FAILED)
	case_value(CL_KERNEL_ARG_INFO_NOT_AVAILABLE)
	case_value(CL_INVALID_VALUE)
	case_value(CL_INVALID_DEVICE_TYPE)
	case_value(CL_INVALID_PLATFORM)
	case_value(CL_INVALID_DEVICE)
	case_value(CL_INVALID_CONTEXT)
	case_value(CL_INVALID_QUEUE_PROPERTIES)
	case_value(CL_INVALID_COMMAND_QUEUE)
	case_value(CL_INVALID_HOST_PTR)
	case_value(CL_INVALID_MEM_OBJECT)
	case_value(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR)
	case_value(CL_INVALID_IMAGE_SIZE)
	case_value(CL_INVALID_SAMPLER)
	case_value(CL_INVALID_BINARY)
	case_value(CL_INVALID_BUILD_OPTIONS)
	case_value(CL_INVALID_PROGRAM)
	case_value(CL_INVALID_PROGRAM_EXECUTABLE)
	case_value(CL_INVALID_KERNEL_NAME)
	case_value(CL_INVALID_KERNEL_DEFINITION)
	case_value(CL_INVALID_KERNEL)
	case_value(CL_INVALID_ARG_INDEX)
	case_value(CL_INVALID_ARG_VALUE)
	case_value(CL_INVALID_ARG_SIZE)
	case_value(CL_INVALID_KERNEL_ARGS)
	case_value(CL_INVALID_WORK_DIMENSION)
	case_value(CL_INVALID_WORK_GROUP_SIZE)
	case_value(CL_INVALID_WORK_ITEM_SIZE)
	case_value(CL_INVALID_GLOBAL_OFFSET)
	case_value(CL_INVALID_EVENT_WAIT_LIST)
	case_value(CL_INVALID_EVENT)
	case_value(CL_INVALID_OPERATION)
	case_value(CL_INVALID_GL_OBJECT)
	case_value(CL_INVALID_BUFFER_SIZE)
	case_value(CL_INVALID_MIP_LEVEL)
	case_value(CL_INVALID_GLOBAL_WORK_SIZE)
	case_value(CL_INVALID_PROPERTY)
	case_value(CL_INVALID_IMAGE_DESCRIPTOR)
	case_value(CL_INVALID_COMPILER_OPTIONS)
	case_value(CL_INVALID_LINKER_OPTIONS)
	case_value(CL_INVALID_DEVICE_PARTITION_COUNT)
#if defined(CL_VERSION_2_0) || defined(CL_VERSION_2_1)
	case_value(CL_INVALID_PIPE_SIZE)
	case_value(CL_INVALID_DEVICE_QUEUE)
#endif
	default:
		return "Unknown error";
	}
}

void cl_utilits::showClError(const cl::Error& e, std::ostream& stream) {
	stream << e.what() << ":" << errorNumberToString(e.err()) << endl;
}

cl::ImageFormat cl_utilits::getFittedImageFormatForGray(const cl::Context& context, cl_mem_flags flags) {
	static cl::ImageFormat matched_formats[] = { { CL_LUMINANCE, CL_UNSIGNED_INT8 }, { CL_LUMINANCE, CL_UNORM_INT8 }, {
			CL_INTENSITY, CL_UNSIGNED_INT8 }, { CL_INTENSITY, CL_UNORM_INT8 }, { CL_A, CL_UNSIGNED_INT8 }, { CL_A,
			CL_UNORM_INT8 }, { CL_R, CL_UNSIGNED_INT8 }, { CL_R, CL_UNORM_INT8 } };
	std::vector<cl::ImageFormat> formats;
	context.getSupportedImageFormats(flags,
	CL_MEM_OBJECT_IMAGE2D, std::addressof(formats));
	for (auto f : matched_formats)
		for (auto itor = formats.begin(), end = formats.end(); itor != end; ++itor)
			if (f.image_channel_data_type == itor->image_channel_data_type
					&& f.image_channel_order == itor->image_channel_order)
				return f;
	throw face_cl_exception("can't found fitted image format  supported by OpenCL device");
}

void cl_utilits::showDevice(const cl::Device& device, std::ostream& stream) {
	stream << "Device Type:" << deviceTypeToString(device.getInfo<CL_DEVICE_TYPE>()) << endl;
	stream << "Device Name:" << std::string(device.getInfo<CL_DEVICE_NAME>().data()) << endl;
	stream << "Vendor:" << std::string(device.getInfo<CL_DEVICE_VENDOR>().data()) << endl;
	stream << "Device Version:" << std::string(device.getInfo<CL_DEVICE_VERSION>().data()) << endl;
	stream << "Device OpenCL C Version:" << std::string(device.getInfo<CL_DEVICE_OPENCL_C_VERSION>().data()) << endl;
	stream << "Driver Version:" << std::string(device.getInfo<CL_DRIVER_VERSION>().data()) << endl;
}

void cl_utilits::showDevices(const std::vector<cl::Device>& devices, std::ostream& stream) {
	for (auto itor = devices.begin(), end = devices.end(); itor != end; ++itor) {
		showDevice(*itor, stream);
		stream << endl;
	}
}

void cl_utilits::showSupportedImageFormats(const cl::Context& context, cl_mem_flags flags,
		cl_mem_object_type type, std::ostream& stream) {
	std::vector<cl::ImageFormat> formats;
	context.getSupportedImageFormats(flags, type, std::addressof(formats));
	stream << "support formats: " << "(" << formats.size() << ") for " << memObjectTypeToString(CL_MEM_OBJECT_IMAGE2D)
			<< " " << memFlagsToString(flags) << endl;
	for (auto itor = formats.begin(), end = formats.end(); itor != end; ++itor) {
		stream << imageFormatToString(*itor) << endl;
	}
}

std::vector<uint8_t> cl_utilits::mem_download(const cl::Memory& memobj) {
	auto ptr = (uint8_t*) memobj.getInfo<CL_MEM_HOST_PTR>();
	auto size = memobj.getInfo<CL_MEM_SIZE>();
	return std::vector<uint8_t>(ptr, ptr + size);
}

std::vector<uint8_t> cl_utilits::image_download(const cl::Image& img) {
	return mem_download(img.getImageInfo<CL_IMAGE_BUFFER>());
}
/* 显示编译日志 */
void cl_utilits::showBuildLog(build_log_type& build_log,const std::string &source_name, std::ostream& stream) {
	stream<<"--------------" << "source name:" << source_name<<"------------------" << endl;
	stream<<"///////////////////////build log begin/////////////////////////////////////"<<endl;
	for (auto itor = build_log.begin(), end = build_log.end(); itor != end; ++itor) {
		showDevice(itor->first, stream);
		stream << std::endl;
		if(itor->second.size()>1)
			stream << itor->second<<endl;
		else
			stream << "build succeed..."<<endl;
	}
	stream<<"///////////////////////build log end/////////////////////////////////////"<<endl;
}
/* 编译一个源码, 参数complie_only指定是否只编译成obj */
cl::Program cl_utilits::buildSource(const build_param& param, bool complie_only) {
	try {
		throw_if_msg(1 != param.sources.size(), "size of build_param::sources must be 1")
#ifdef CL_VERSION_1_2
		ProgramExt program(param.context, param.sources[0].second);
#else
		cl::Program program(param.context, param.sources[0].second);
#endif
		try {
#ifdef CL_VERSION_1_2
			if (complie_only) {
				show_on_buildstart(_DEF_STRING(complie), param.compiler_options, param.sources[0]);
				program.compile(param.devices,param.input_headers,param.compiler_options.data());
			} else
#else
			throw_exception_if(face_cl_build_exception,complie_only,"unsupported complie under version OpenCL 1.2")
#endif
			{
				auto options = param.build_options();
				show_on_buildstart(_DEF_STRING(build), options, param.sources[0]);
				if (param.devices.size())
					program.build(param.devices, options.data());
				else
					program.build(options.data());
			}
			show_on_buildend(program, param.sources[0]);
			return std::move(program);
		}
#ifdef CL_VERSION_2_0
		//当OpenCL版本为2.0以上时，编译错误抛出cl::BuildError异常
		catch (cl::BuildError &e) {
			auto log = e.getBuildLog();
			showBuildLog(log, param.sources[0].first);
			throw face_cl_build_exception(SOURCE_AT, e);
		}
#else
		//当OpenCL版本为1.1,1.2时，编译错误抛出cl::Error异常
		catch (cl::Error& e) {
			auto log = cl_utilits::getBuildInfo<CL_PROGRAM_BUILD_LOG>(program);
			showBuildLog(log, param.sources[0].first);
			throw face_cl_build_exception(SOURCE_AT, e, log);
		}
#endif
	} catch (cl::Error& e) {
		throw face_cl_build_exception(SOURCE_AT, e);
	}
}
#if defined(CL_VERSION_2_0) && defined(CL_HPP_ENABLE_EXCEPTIONS) // 判断 #include <CL/cl2.hpp>
/* 封装cl::linkProgram 出错抛出face_cl_build_exception */
static  cl::Program cl_utilits::linkProgram(
	std::vector<cl::Program> inputPrograms,
	const char* options,
	void (CL_CALLBACK * notifyFptr)(cl_program, void *),
	void* data,
	cl_int* err)
{
	try{
		show_on_buildstart(_DEF_STRING(link),options,{Unknow_Name,Unknow_Name});
		return cl::linkProgram(inputPrograms,options,notifyFptr,data,err);
	}catch(cl::BuildError&e){
		throw face_cl_build_exception(SOURCE_AT,e);
	}
}
#elif defined(CL_VERSION_1_2)
/* 封装cl::linkProgram 出错抛出face_cl_build_exception */
cl::Program cl_utilits::linkProgram(
	std::vector<cl::Program> inputPrograms,
	const char* compiler_options,
	void (CL_CALLBACK * notifyFptr)(cl_program, void *),
	void* data,
	cl_int* err)
{
	cl_int error_local = CL_SUCCESS;

	cl_program * programs = (cl_program*)alloca(inputPrograms.size() * sizeof(cl_program));

	if (programs != NULL) {
		for (unsigned int i = 0; i < inputPrograms.size(); i++) {
			programs[i] = inputPrograms[i]();
		}
	}

	cl::Context ctx;
	if (inputPrograms.size() > 0) {
		ctx = inputPrograms[0].getInfo<CL_PROGRAM_CONTEXT>(&error_local);
		if (error_local != CL_SUCCESS) {
			cl::detail::errHandler(error_local, _DEF_STRING(clLinkProgram));
		}
	}
	cl_program prog = ::clLinkProgram(
		ctx(),
		0,
		NULL,
		compiler_options,
		(cl_uint)inputPrograms.size(),
		programs,
		notifyFptr,
		data,
		&error_local);
	cl::Program program(prog);
	if (error_local != CL_SUCCESS) {
		try{
			auto log = cl_utilits::getBuildInfo<CL_PROGRAM_BUILD_LOG>(program);
			showBuildLog(log);
			throw face_cl_build_exception(SOURCE_AT, cl::Error(error_local, _DEF_STRING(clCompileProgram)), log);
		}catch (cl::Error &) {
			// cl_utilits::getBuildInfo获取log失败时只输出cl::Error
			throw face_cl_build_exception(SOURCE_AT, cl::Error(error_local, _DEF_STRING(clCompileProgram)));
		}
	}
	return std::move(program);
}
#endif

#ifdef CL_VERSION_1_2
cl::Program cl_utilits::compileSource(const build_param& param){
	return buildSource(param,true);
}
/* 编译连接多个源码生成可执行程序  */
cl::Program cl_utilits::buildMultiFilesProgram(const build_param& param) {
	throw_if_msg(0==param.sources.size(),"size of build_param::sources must not be 0")
	std::vector<cl::Program> objs;
	auto p=param;
	for (auto source : param.sources) {
		p.sources = { source };
		objs.emplace_back(compileSource(p));
	}
	return linkProgram(objs, param.linker_options.data()); // 连接obj生成可执行程序
}
#endif
/* (通用)编译源码源码生成可执行程序  */
cl::Program cl_utilits::buildExecutableProgram(const build_param& param){
	// 如果指定了input_headers，则必须complie/link
	if(1==param.sources.size()&&0==param.input_headers.size()){
		return buildSource(param	);
	}
#ifdef CL_VERSION_1_2
	return buildMultiFilesProgram(param);
#else
	throw face_cl_exception("multifile complie can't be supported on OpenCL under 1.2 version");
#endif
}

#if defined(CL_VERSION_1_2)
cl_int ProgramExt::compile(
	const std::vector<cl::Device> &devices,
	const std::vector<program_info_type> &input_headers,
	const char* options,
	void (CL_CALLBACK * notifyFptr)(cl_program, void *),
	void* data) const
{
	// 生成cl_device_id数组
	auto deviceIDs=cl::cl_c_vector(devices);
	// 生成cl_program数组
	auto headers=cl::cl_c_vector2(input_headers);
	// 生成头文件名字数组
	auto include_names=cl::cl_c_vector1(input_headers);
	return cl::detail::errHandler(
		::clCompileProgram(
			object_,
			(cl_uint)deviceIDs.size(),
			deviceIDs.size()?deviceIDs.data():nullptr,
			options,
			(cl_uint)headers.size(),
			headers.size()?headers.data():nullptr,
			(const char**)(headers.size()?include_names.data():nullptr),
			notifyFptr,
			data),
			_DEF_STRING(clCompileProgram));
}
#endif

/* 通过param提供的源码创建一组cl::Kernel,并将cl::Kernel以名字为key加入kernels映射表中  */
cl_utilits::kernel_table_type cl_utilits::createKernels(const build_param& param) {
	auto program = buildExecutableProgram(param);
	kernel_table_type map;
	std::vector<cl::Kernel> kernels;
	program.createKernels(std::addressof(kernels));
	for (auto k : kernels) {
		auto name = k.getInfo<CL_KERNEL_FUNCTION_NAME>();
		map.emplace(std::string(name.data()), k);
#ifndef NDEBUG
		cout << "kernal function name:" << name << endl;
#endif
	}
	return std::move(map);
}

cl::Buffer cl_utilits::createBuffer(cl_mem_flags flags,const cl::Context &context,size_t size,const void* ptr){
	// flag为CL_MEM_COPY_HOST_PTR|CL_MEM_USE_HOST_PTR时,数组不能为空
	throw_if((flags&(CL_MEM_COPY_HOST_PTR|CL_MEM_USE_HOST_PTR))&&(0==size||nullptr==ptr))
	return cl::Buffer (
			context,
			flags,
			size,
			(void*)(0==size?nullptr:ptr));
}

cl::Image2D cl_utilits::createImage2D(cl_mem_flags flags,const cl::Context &context,
		const cl::ImageFormat& format,
		size_t width,
		size_t height,
		size_t row_pitch,
		const void* ptr){
	throw_if_msg(!context(),"uninitialized cl::Memory object")
	throw_if((flags&(CL_MEM_COPY_HOST_PTR|CL_MEM_USE_HOST_PTR))&&(nullptr==ptr))
	return cl::Image2D (
			context,
			flags,
			format,
			width,
			height,
			nullptr==ptr?0:row_pitch,// This must be 0 if host_ptr is NULL
			(void*)ptr);
}
/*
 * 获取OpenCL设备每个计算单元(CU)中处理单元(PE)个数
 */
std::pair<size_t, size_t> cl_utilits::kernel_preferred_work_group_size_multiple(const cl::Device& device)  {
	cl::Program program(cl::Context({device}), "__kernel void test(){}", true);
	std::vector<cl::Kernel> kernels;
	program.createKernels(std::addressof(kernels));
	auto kpwgsm = kernels[0].getWorkGroupInfo<CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE>(device);
	auto shift_num = _bsr_int32_(uint32_t(kpwgsm));
	return std::make_pair(kpwgsm, shift_num);
}

#ifdef CL_FILES_EMBEDDED
std::string load_string_from_code(const std::string &path);
#define _LOAD_PROGRAM_STRING(s)  load_string_from_code(s)
#else
#define _LOAD_PROGRAM_STRING(s)  load_string(s.data())
#endif
std::string load_program_string(const std::string &path){
	return _LOAD_PROGRAM_STRING(path);
}
#undef _LOAD_PROGRAM_STRING
const std::string cl_utilits::Unknow_Name="unknow_name";
const std::string cl_utilits::Empty_String="";

} /* namespace gdface */
